/**
 * @file test_synchronization.cpp
 * @author 心流 (iFlow CLI)
 * @brief 测试OpenCL同步原语功能
 * @version 0.1
 * @date 2025-11-27
 *
 * @copyright Copyright (c) 2025
 *
 */
#define BOOST_TEST_MODULE TestSynchronization
#include <CL/cl.h>

#include <boost/format.hpp>
#include <boost/test/included/unit_test.hpp>
#include <boost/timer/timer.hpp>
#include <chrono>
#include <thread>
#include <vector>

#include "test_fixture.hpp"
#include "utility.hpp"

// 测试clFinish函数
BOOST_AUTO_TEST_CASE(TestFinish)
{
    CLTestFixture    fixture;
    cl_context       context = fixture.context_;
    cl_command_queue queue   = fixture.queue_;
    cl_device_id     device  = fixture.device_;
    BOOST_CHECK(context != nullptr);
    BOOST_CHECK(queue != nullptr);
    BOOST_CHECK(device != nullptr);

    // 创建简单的内核源代码
    const char* kernel_source = R"(
        __kernel void test_kernel(__global int* data) {
            int id = get_global_id(0);
            data[id] = id * id;
        }
    )";

    // 创建程序
    cl_program program =
        clCreateProgramWithSource(context, 1, &kernel_source, nullptr, nullptr);
    BOOST_CHECK(program != nullptr);

    // 构建程序
    cl_int err = clBuildProgram(program, 1, &device, nullptr, nullptr, nullptr);
    if (err != CL_SUCCESS)
    {
        size_t log_size;
        clGetProgramBuildInfo(
            program, device, CL_PROGRAM_BUILD_LOG, 0, nullptr, &log_size);
        std::vector<char> log(log_size);
        clGetProgramBuildInfo(program,
                              device,
                              CL_PROGRAM_BUILD_LOG,
                              log_size,
                              log.data(),
                              nullptr);
        std::cout << "Build log: " << log.data() << std::endl;
    }
    BOOST_CHECK_EQUAL(err, CL_SUCCESS);

    // 创建内核
    cl_kernel kernel = clCreateKernel(program, "test_kernel", &err);
    BOOST_CHECK_EQUAL(err, CL_SUCCESS);
    BOOST_CHECK(kernel != nullptr);

    // 创建内存对象
    const size_t buffer_size = 1024 * sizeof(int);
    cl_mem       buffer =
        clCreateBuffer(context, CL_MEM_READ_WRITE, buffer_size, nullptr, &err);
    BOOST_CHECK_EQUAL(err, CL_SUCCESS);
    BOOST_CHECK(buffer != nullptr);

    // 设置内核参数
    err = clSetKernelArg(kernel, 0, sizeof(cl_mem), &buffer);
    BOOST_CHECK_EQUAL(err, CL_SUCCESS);

    // 执行内核（异步）
    size_t   global_work_size = 1024;
    cl_event kernel_event     = nullptr;
    err                       = clEnqueueNDRangeKernel(queue,
                                 kernel,
                                 1,
                                 nullptr,
                                 &global_work_size,
                                 nullptr,
                                 0,
                                 nullptr,
                                 &kernel_event);
    BOOST_CHECK_EQUAL(err, CL_SUCCESS);
    BOOST_CHECK(kernel_event != nullptr);

    // 调用clFinish等待所有命令完成
    err = clFinish(queue);
    BOOST_CHECK_EQUAL(err, CL_SUCCESS);

    // 确认事件已完成
    cl_int event_status = 0;
    err                 = clGetEventInfo(kernel_event,
                         CL_EVENT_COMMAND_EXECUTION_STATUS,
                         sizeof(event_status),
                         &event_status,
                         nullptr);
    BOOST_CHECK_EQUAL(err, CL_SUCCESS);
    BOOST_CHECK_EQUAL(event_status, CL_COMPLETE);

    // 清理资源
    err = clReleaseEvent(kernel_event);
    BOOST_CHECK_EQUAL(err, CL_SUCCESS);

    err = clReleaseMemObject(buffer);
    BOOST_CHECK_EQUAL(err, CL_SUCCESS);

    err = clReleaseKernel(kernel);
    BOOST_CHECK_EQUAL(err, CL_SUCCESS);

    err = clReleaseProgram(program);
    BOOST_CHECK_EQUAL(err, CL_SUCCESS);
}

// 测试clFlush函数
BOOST_AUTO_TEST_CASE(TestFlush)
{
    CLWrapper cl;
    bool      init_result = cl.Init();
    BOOST_CHECK(init_result == true);

    cl_context       context = cl.GetContext();
    cl_command_queue queue   = cl.GetQueue();
    cl_device_id     device  = cl.GetDeviceID();
    BOOST_CHECK(context != nullptr);
    BOOST_CHECK(queue != nullptr);
    BOOST_CHECK(device != nullptr);

    // 创建简单的内核源代码
    const char* kernel_source = R"(
        __kernel void test_kernel(__global int* data) {
            int id = get_global_id(0);
            data[id] = id + 1;
        }
    )";

    // 创建程序
    cl_program program =
        clCreateProgramWithSource(context, 1, &kernel_source, nullptr, nullptr);
    BOOST_CHECK(program != nullptr);

    // 构建程序
    cl_int err = clBuildProgram(program, 1, &device, nullptr, nullptr, nullptr);
    if (err != CL_SUCCESS)
    {
        size_t log_size;
        clGetProgramBuildInfo(
            program, device, CL_PROGRAM_BUILD_LOG, 0, nullptr, &log_size);
        std::vector<char> log(log_size);
        clGetProgramBuildInfo(program,
                              device,
                              CL_PROGRAM_BUILD_LOG,
                              log_size,
                              log.data(),
                              nullptr);
        std::cout << "Build log: " << log.data() << std::endl;
    }
    BOOST_CHECK_EQUAL(err, CL_SUCCESS);

    // 创建内核
    cl_kernel kernel = clCreateKernel(program, "test_kernel", &err);
    BOOST_CHECK_EQUAL(err, CL_SUCCESS);
    BOOST_CHECK(kernel != nullptr);

    // 创建内存对象
    const size_t buffer_size = 512 * sizeof(int);
    cl_mem       buffer =
        clCreateBuffer(context, CL_MEM_READ_WRITE, buffer_size, nullptr, &err);
    BOOST_CHECK_EQUAL(err, CL_SUCCESS);
    BOOST_CHECK(buffer != nullptr);

    // 设置内核参数
    err = clSetKernelArg(kernel, 0, sizeof(cl_mem), &buffer);
    BOOST_CHECK_EQUAL(err, CL_SUCCESS);

    // 执行内核（异步）
    size_t   global_work_size = 512;
    cl_event kernel_event     = nullptr;
    err                       = clEnqueueNDRangeKernel(queue,
                                 kernel,
                                 1,
                                 nullptr,
                                 &global_work_size,
                                 nullptr,
                                 0,
                                 nullptr,
                                 &kernel_event);
    BOOST_CHECK_EQUAL(err, CL_SUCCESS);
    BOOST_CHECK(kernel_event != nullptr);

    // 调用clFlush刷新命令队列
    err = clFlush(queue);
    BOOST_CHECK_EQUAL(err, CL_SUCCESS);

    // 验证命令已经提交（但不一定执行完成）
    cl_int event_status = 0;
    err                 = clGetEventInfo(kernel_event,
                         CL_EVENT_COMMAND_EXECUTION_STATUS,
                         sizeof(event_status),
                         &event_status,
                         nullptr);
    BOOST_CHECK_EQUAL(err, CL_SUCCESS);
    // 事件应该至少已经提交
    BOOST_CHECK(event_status <= CL_SUBMITTED);

    // 等待事件完成
    err = clWaitForEvents(1, &kernel_event);
    BOOST_CHECK_EQUAL(err, CL_SUCCESS);

    // 确认事件已完成
    err = clGetEventInfo(kernel_event,
                         CL_EVENT_COMMAND_EXECUTION_STATUS,
                         sizeof(event_status),
                         &event_status,
                         nullptr);
    BOOST_CHECK_EQUAL(err, CL_SUCCESS);
    BOOST_CHECK_EQUAL(event_status, CL_COMPLETE);

    // 清理资源
    err = clReleaseEvent(kernel_event);
    BOOST_CHECK_EQUAL(err, CL_SUCCESS);

    err = clReleaseMemObject(buffer);
    BOOST_CHECK_EQUAL(err, CL_SUCCESS);

    err = clReleaseKernel(kernel);
    BOOST_CHECK_EQUAL(err, CL_SUCCESS);

    err = clReleaseProgram(program);
    BOOST_CHECK_EQUAL(err, CL_SUCCESS);
}

// 测试多个命令和队列同步
BOOST_AUTO_TEST_CASE(TestMultipleCommandsSync)
{
    CLWrapper cl;
    bool      init_result = cl.Init();
    BOOST_CHECK(init_result == true);

    cl_context       context = cl.GetContext();
    cl_command_queue queue   = cl.GetQueue();
    cl_device_id     device  = cl.GetDeviceID();
    BOOST_CHECK(context != nullptr);
    BOOST_CHECK(queue != nullptr);
    BOOST_CHECK(device != nullptr);

    // 创建内核源代码
    const char* kernel_source = R"(
        __kernel void increment_kernel(__global int* data) {
            int id = get_global_id(0);
            data[id] = data[id] + 1;
        }
    )";

    // 创建程序
    cl_program program =
        clCreateProgramWithSource(context, 1, &kernel_source, nullptr, nullptr);
    BOOST_CHECK(program != nullptr);

    // 构建程序
    cl_int err = clBuildProgram(program, 1, &device, nullptr, nullptr, nullptr);
    if (err != CL_SUCCESS)
    {
        size_t log_size;
        clGetProgramBuildInfo(
            program, device, CL_PROGRAM_BUILD_LOG, 0, nullptr, &log_size);
        std::vector<char> log(log_size);
        clGetProgramBuildInfo(program,
                              device,
                              CL_PROGRAM_BUILD_LOG,
                              log_size,
                              log.data(),
                              nullptr);
        std::cout << "Build log: " << log.data() << std::endl;
    }
    BOOST_CHECK_EQUAL(err, CL_SUCCESS);

    // 创建内核
    cl_kernel kernel = clCreateKernel(program, "increment_kernel", &err);
    BOOST_CHECK_EQUAL(err, CL_SUCCESS);
    BOOST_CHECK(kernel != nullptr);

    // 创建内存对象
    const size_t buffer_size = 256 * sizeof(int);
    cl_mem       buffer =
        clCreateBuffer(context, CL_MEM_READ_WRITE, buffer_size, nullptr, &err);
    BOOST_CHECK_EQUAL(err, CL_SUCCESS);
    BOOST_CHECK(buffer != nullptr);

    // 初始化数据
    std::vector<int> initial_data(256, 0);
    err = clEnqueueWriteBuffer(queue,
                               buffer,
                               CL_TRUE,
                               0,
                               buffer_size,
                               initial_data.data(),
                               0,
                               nullptr,
                               nullptr);
    BOOST_CHECK_EQUAL(err, CL_SUCCESS);

    // 设置内核参数
    err = clSetKernelArg(kernel, 0, sizeof(cl_mem), &buffer);
    BOOST_CHECK_EQUAL(err, CL_SUCCESS);

    // 执行多个内核执行命令
    size_t                global_work_size = 256;
    std::vector<cl_event> events(5);

    for (int i = 0; i < 5; i++)
    {
        err = clEnqueueNDRangeKernel(queue,
                                     kernel,
                                     1,
                                     nullptr,
                                     &global_work_size,
                                     nullptr,
                                     0,
                                     nullptr,
                                     &events[i]);
        BOOST_CHECK_EQUAL(err, CL_SUCCESS);
        BOOST_CHECK(events[i] != nullptr);
    }

    // 刷新队列
    err = clFlush(queue);
    BOOST_CHECK_EQUAL(err, CL_SUCCESS);

    // 验证所有事件至少已提交
    for (int i = 0; i < 5; i++)
    {
        cl_int event_status = 0;
        err                 = clGetEventInfo(events[i],
                             CL_EVENT_COMMAND_EXECUTION_STATUS,
                             sizeof(event_status),
                             &event_status,
                             nullptr);
        BOOST_CHECK_EQUAL(err, CL_SUCCESS);
        BOOST_CHECK(event_status <= CL_SUBMITTED);
    }

    // 使用clFinish等待所有命令完成
    err = clFinish(queue);
    BOOST_CHECK_EQUAL(err, CL_SUCCESS);

    // 验证所有事件都已完成
    for (int i = 0; i < 5; i++)
    {
        cl_int event_status = 0;
        err                 = clGetEventInfo(events[i],
                             CL_EVENT_COMMAND_EXECUTION_STATUS,
                             sizeof(event_status),
                             &event_status,
                             nullptr);
        BOOST_CHECK_EQUAL(err, CL_SUCCESS);
        BOOST_CHECK_EQUAL(event_status, CL_COMPLETE);
    }

    // 读取结果验证
    std::vector<int> result_data(256);
    err = clEnqueueReadBuffer(queue,
                              buffer,
                              CL_TRUE,
                              0,
                              buffer_size,
                              result_data.data(),
                              0,
                              nullptr,
                              nullptr);
    BOOST_CHECK_EQUAL(err, CL_SUCCESS);

    // 验证数据被正确增加了5次
    for (int i = 0; i < 256; i++)
    {
        BOOST_CHECK_EQUAL(result_data[i], 5);  // 初始值0，执行了5次+1操作
    }

    // 清理资源
    for (int i = 0; i < 5; i++)
    {
        err = clReleaseEvent(events[i]);
        BOOST_CHECK_EQUAL(err, CL_SUCCESS);
    }

    err = clReleaseMemObject(buffer);
    BOOST_CHECK_EQUAL(err, CL_SUCCESS);

    err = clReleaseKernel(kernel);
    BOOST_CHECK_EQUAL(err, CL_SUCCESS);

    err = clReleaseProgram(program);
    BOOST_CHECK_EQUAL(err, CL_SUCCESS);
}

// 测试同步函数错误处理
BOOST_AUTO_TEST_CASE(TestSynchronizationErrorHandling)
{
    CLWrapper cl;
    bool      init_result = cl.Init();
    BOOST_CHECK(init_result == true);

    cl_command_queue queue = cl.GetQueue();
    BOOST_CHECK(queue != nullptr);

    cl_int err = CL_SUCCESS;

    // 测试使用无效队列调用clFinish（应该失败）
    err = clFinish(nullptr);
    BOOST_CHECK_NE(err, CL_SUCCESS);  // 应该失败

    // 测试使用无效队列调用clFlush（应该失败）
    err = clFlush(nullptr);
    BOOST_CHECK_NE(err, CL_SUCCESS);  // 应该失败

    // 正常调用clFlush和clFinish
    err = clFlush(queue);
    BOOST_CHECK_EQUAL(err, CL_SUCCESS);

    err = clFinish(queue);
    BOOST_CHECK_EQUAL(err, CL_SUCCESS);
}

// 测试clFinish和clFlush性能对比
BOOST_AUTO_TEST_CASE(TestSyncPerformance)
{
    CLWrapper cl;
    bool      init_result = cl.Init();
    BOOST_CHECK(init_result == true);

    cl_context       context = cl.GetContext();
    cl_command_queue queue   = cl.GetQueue();
    cl_device_id     device  = cl.GetDeviceID();
    BOOST_CHECK(context != nullptr);
    BOOST_CHECK(queue != nullptr);
    BOOST_CHECK(device != nullptr);

    // 创建内核源代码
    const char* kernel_source = R"(
        __kernel void delay_kernel(__global int* data) {
            int id = get_global_id(0);
            int sum = 0;
            for (int i = 0; i < 1000; i++) {
                sum += i;
            }
            data[id] = sum;
        }
    )";

    // 创建程序
    cl_program program =
        clCreateProgramWithSource(context, 1, &kernel_source, nullptr, nullptr);
    BOOST_CHECK(program != nullptr);

    // 构建程序
    cl_int err = clBuildProgram(program, 1, &device, nullptr, nullptr, nullptr);
    if (err != CL_SUCCESS)
    {
        size_t log_size;
        clGetProgramBuildInfo(
            program, device, CL_PROGRAM_BUILD_LOG, 0, nullptr, &log_size);
        std::vector<char> log(log_size);
        clGetProgramBuildInfo(program,
                              device,
                              CL_PROGRAM_BUILD_LOG,
                              log_size,
                              log.data(),
                              nullptr);
        std::cout << "Build log: " << log.data() << std::endl;
    }
    BOOST_CHECK_EQUAL(err, CL_SUCCESS);

    // 创建内核
    cl_kernel kernel = clCreateKernel(program, "delay_kernel", &err);
    BOOST_CHECK_EQUAL(err, CL_SUCCESS);
    BOOST_CHECK(kernel != nullptr);

    // 创建内存对象
    const size_t buffer_size = 64 * sizeof(int);
    cl_mem       buffer =
        clCreateBuffer(context, CL_MEM_READ_WRITE, buffer_size, nullptr, &err);
    BOOST_CHECK_EQUAL(err, CL_SUCCESS);
    BOOST_CHECK(buffer != nullptr);

    // 设置内核参数
    err = clSetKernelArg(kernel, 0, sizeof(cl_mem), &buffer);
    BOOST_CHECK_EQUAL(err, CL_SUCCESS);

    // 测试clFlush - 应该快速返回
    size_t   global_work_size = 64;
    cl_event kernel_event     = nullptr;
    err                       = clEnqueueNDRangeKernel(queue,
                                 kernel,
                                 1,
                                 nullptr,
                                 &global_work_size,
                                 nullptr,
                                 0,
                                 nullptr,
                                 &kernel_event);
    BOOST_CHECK_EQUAL(err, CL_SUCCESS);
    BOOST_CHECK(kernel_event != nullptr);

    // 记录clFlush开始时间
    auto flush_start = std::chrono::high_resolution_clock::now();
    err              = clFlush(queue);
    auto flush_end   = std::chrono::high_resolution_clock::now();
    BOOST_CHECK_EQUAL(err, CL_SUCCESS);

    // clFlush应该几乎立即返回
    auto flush_duration = std::chrono::duration_cast<std::chrono::milliseconds>(
        flush_end - flush_start);
    // clFlush应该在很短的时间内返回（比如10ms内）
    // 这个检查可能在某些系统上失败，但作为性能验证是有用的

    // 记录clFinish开始时间
    auto finish_start = std::chrono::high_resolution_clock::now();
    err               = clFinish(queue);
    auto finish_end   = std::chrono::high_resolution_clock::now();
    BOOST_CHECK_EQUAL(err, CL_SUCCESS);

    // clFinish应该需要更长时间，因为它等待命令完成
    auto finish_duration =
        std::chrono::duration_cast<std::chrono::milliseconds>(finish_end -
                                                              finish_start);

    // 等待事件完成
    err = clWaitForEvents(1, &kernel_event);
    BOOST_CHECK_EQUAL(err, CL_SUCCESS);

    // 清理资源
    err = clReleaseEvent(kernel_event);
    BOOST_CHECK_EQUAL(err, CL_SUCCESS);

    err = clReleaseMemObject(buffer);
    BOOST_CHECK_EQUAL(err, CL_SUCCESS);

    err = clReleaseKernel(kernel);
    BOOST_CHECK_EQUAL(err, CL_SUCCESS);

    err = clReleaseProgram(program);
    BOOST_CHECK_EQUAL(err, CL_SUCCESS);
}

// 测试内存屏障和队列同步
BOOST_AUTO_TEST_CASE(TestMemoryBarrierSync)
{
    CLWrapper cl;
    bool      init_result = cl.Init();
    BOOST_CHECK(init_result == true);

    cl_context       context = cl.GetContext();
    cl_command_queue queue   = cl.GetQueue();
    cl_device_id     device  = cl.GetDeviceID();
    BOOST_CHECK(context != nullptr);
    BOOST_CHECK(queue != nullptr);
    BOOST_CHECK(device != nullptr);

    // 创建内核源代码，包含内存屏障测试
    const char* kernel_source = R"(
        __kernel void barrier_test_kernel(__global int* input, __global int* output) {
            int id = get_global_id(0);
            output[id] = input[id] * 2;
        }
    )";

    // 创建程序
    cl_program program =
        clCreateProgramWithSource(context, 1, &kernel_source, nullptr, nullptr);
    BOOST_CHECK(program != nullptr);

    // 构建程序
    cl_int err = clBuildProgram(program, 1, &device, nullptr, nullptr, nullptr);
    if (err != CL_SUCCESS)
    {
        size_t log_size;
        clGetProgramBuildInfo(
            program, device, CL_PROGRAM_BUILD_LOG, 0, nullptr, &log_size);
        std::vector<char> log(log_size);
        clGetProgramBuildInfo(program,
                              device,
                              CL_PROGRAM_BUILD_LOG,
                              log_size,
                              log.data(),
                              nullptr);
        std::cout << "Build log: " << log.data() << std::endl;
    }
    BOOST_CHECK_EQUAL(err, CL_SUCCESS);

    // 创建内核
    cl_kernel kernel = clCreateKernel(program, "barrier_test_kernel", &err);
    BOOST_CHECK_EQUAL(err, CL_SUCCESS);
    BOOST_CHECK(kernel != nullptr);

    // 创建内存对象
    const size_t buffer_size = 128 * sizeof(int);
    cl_mem       input_buffer =
        clCreateBuffer(context, CL_MEM_READ_WRITE, buffer_size, nullptr, &err);
    BOOST_CHECK_EQUAL(err, CL_SUCCESS);
    BOOST_CHECK(input_buffer != nullptr);

    cl_mem output_buffer =
        clCreateBuffer(context, CL_MEM_READ_WRITE, buffer_size, nullptr, &err);
    BOOST_CHECK_EQUAL(err, CL_SUCCESS);
    BOOST_CHECK(output_buffer != nullptr);

    // 初始化输入数据
    std::vector<int> input_data(128);
    for (int i = 0; i < 128; i++)
    {
        input_data[i] = i;
    }

    // 异步写入输入数据
    cl_event write_event = nullptr;
    err                  = clEnqueueWriteBuffer(queue,
                               input_buffer,
                               CL_FALSE,
                               0,
                               buffer_size,
                               input_data.data(),
                               0,
                               nullptr,
                               &write_event);
    BOOST_CHECK_EQUAL(err, CL_SUCCESS);
    BOOST_CHECK(write_event != nullptr);

    // 设置内核参数
    err = clSetKernelArg(kernel, 0, sizeof(cl_mem), &input_buffer);
    BOOST_CHECK_EQUAL(err, CL_SUCCESS);
    err = clSetKernelArg(kernel, 1, sizeof(cl_mem), &output_buffer);
    BOOST_CHECK_EQUAL(err, CL_SUCCESS);

    // 执行内核，依赖写入事件
    size_t   global_work_size = 128;
    cl_event kernel_event     = nullptr;
    err                       = clEnqueueNDRangeKernel(queue,
                                 kernel,
                                 1,
                                 nullptr,
                                 &global_work_size,
                                 nullptr,
                                 1,
                                 &write_event,
                                 &kernel_event);
    BOOST_CHECK_EQUAL(err, CL_SUCCESS);
    BOOST_CHECK(kernel_event != nullptr);

    // 异步读取输出数据，依赖内核事件
    std::vector<int> output_data(128);
    cl_event         read_event = nullptr;
    err                         = clEnqueueReadBuffer(queue,
                              output_buffer,
                              CL_FALSE,
                              0,
                              buffer_size,
                              output_data.data(),
                              1,
                              &kernel_event,
                              &read_event);
    BOOST_CHECK_EQUAL(err, CL_SUCCESS);
    BOOST_CHECK(read_event != nullptr);

    // 使用clFinish确保所有操作完成
    err = clFinish(queue);
    BOOST_CHECK_EQUAL(err, CL_SUCCESS);

    // 验证结果
    for (int i = 0; i < 128; i++)
    {
        BOOST_CHECK_EQUAL(output_data[i], input_data[i] * 2);
    }

    // 清理资源
    err = clReleaseEvent(write_event);
    BOOST_CHECK_EQUAL(err, CL_SUCCESS);

    err = clReleaseEvent(kernel_event);
    BOOST_CHECK_EQUAL(err, CL_SUCCESS);

    err = clReleaseEvent(read_event);
    BOOST_CHECK_EQUAL(err, CL_SUCCESS);

    err = clReleaseMemObject(input_buffer);
    BOOST_CHECK_EQUAL(err, CL_SUCCESS);

    err = clReleaseMemObject(output_buffer);
    BOOST_CHECK_EQUAL(err, CL_SUCCESS);

    err = clReleaseKernel(kernel);
    BOOST_CHECK_EQUAL(err, CL_SUCCESS);

    err = clReleaseProgram(program);
    BOOST_CHECK_EQUAL(err, CL_SUCCESS);
}